#include "common.h"
#include "kernel_operator.h"
#include <cstdint>
namespace AscendC
{
template <typename _DT_X, typename _DT_CLIP_VALUE_MIN, typename _DT_CLIP_VALUE_MAX, typename _DT_Y> class Kernel
{
  public:
    TPipe pipe;
    CommonTiling tilingData;
    _DT_CLIP_VALUE_MIN min_value;
    _DT_CLIP_VALUE_MAX max_value;
    DefInTensor(X);
    DefOutTensor(Y);

  public:
    __aicore__ inline Kernel()
    {
    }
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR clip_value_min, GM_ADDR clip_value_max, GM_ADDR y,
                                CommonTiling &tiling)
    {
        tilingData = tiling;
        // 切分global
        SetGlobalBuffer(X, x, tilingData.startIdx, tilingData.blockLength);
        SetGlobalBuffer(Y, y, tilingData.startIdx, tilingData.blockLength);
        // 初始化队列
        InitQueueSimple(X, tilingData.bufferSize);
        InitQueueSimple(Y, tilingData.bufferSize);
        // 初始化buf
        // CLIP_VALUE_MIN = min_value
        min_value = CAST_GM_ADDR(TypeOf(CLIP_VALUE_MIN), clip_value_min, 0)[0];
        // CLIP_VALUE_MAX = max_value
        max_value = CAST_GM_ADDR(TypeOf(CLIP_VALUE_MAX), clip_value_max, 0)[0];
    }
    __aicore__ inline void Process()
    {
        auto finnal_progress = tilingData.loopCount - 1;
        for (uint32_t i = 0; i < tilingData.loopCount - 1; ++i)
        {
            CopyIn(i, tilingData.bufferSize);
            Compute(i, tilingData.bufferSize);
            CopyOut(i, tilingData.bufferSize);
        }
        auto finalCopyCount = tilingData.getFinalCopyCount<TypeOf(X)>();
        CopyIn(finnal_progress, finalCopyCount);
        Compute(finnal_progress, tilingData.finalCalcCount);
        CopyOut(finnal_progress, finalCopyCount);
    }
    __aicore__ inline void CopyIn(uint32_t progress, uint32_t calcCount)
    {
        auto index = progress * tilingData.bufferSize;
        EnQueGlobal2Local(X, index, calcCount);
    }
    __aicore__ inline void Compute(uint32_t progress, uint32_t calcCount)
    {
        DeQueSimple(X);
        QueAllocSimple(Y);
        Mins(LTensorName(X), LTensorName(X), max_value, calcCount);
        Maxs(LTensorName(Y), LTensorName(X), min_value, calcCount);
        EnQue(Y);
        QueFree(X);
    }
    __aicore__ inline void CopyOut(uint32_t progress, uint32_t calcCount)
    {
        auto index = progress * tilingData.bufferSize;
        DeQueLocal2Global(Y, index, calcCount)
    }
};
} // namespace AscendC
extern "C" __global__ __aicore__ void clip_by_value(GM_ADDR x, GM_ADDR clip_value_min, GM_ADDR clip_value_max,
                                                    GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
{
    // 获取tilingData
    GET_TILING_DATA(tiling_data, tiling);
    CRATE_COMMON_TILING_SIMPLE;
    AscendC::Kernel<DTYPE_X, DTYPE_CLIP_VALUE_MIN, DTYPE_CLIP_VALUE_MAX, DTYPE_Y> op;
    op.Init(x, clip_value_min, clip_value_max, y, commonTiling);
    op.Process();
}